/*
============================================================================
Name        : GameRunner.cu
Author      : sagnitude <leyicong@gmail.com>
Version     :
Copyright   : Copyright All Rights Reserved.
Description : CUDA version for Othello Game implementation
============================================================================
*/

#include <iostream>
#include <fstream>
#include <sstream>
#include <time.h>
#include <numeric>
#include <stdlib.h>
#include <curand.h>
#include "inc.h"


#define SYNCED(x) {\
    pthread_mutex_lock( &sum_mutex );\
    x;\
    pthread_mutex_unlock( &sum_mutex );\
}

#define DEBUG 0
#define LOG(x) if (DEBUG) {x;}

#define RECORD 1
#define REC(x) if (RECORD) SYNCED(x)

#define LOGRESULT 0
#define LOGR(x) if (!DEBUG & LOGRESULT) {x;}

#define MATRIX_WIDTH 8
#define DX(x) directionX[x]
#define DY(x) directionY[x]
#define FV(x) factionValues[x]
#define FOFF(x) (factionOffset[x])

#define GET_BIT_AT(x, y) (!!(x & (1 << (31 - y))))
#define SET_BIT_AT(x, v, n) ;(x = (x | (1 << (31 - n))) & ((v << (31 - n)) | ((~0) ^ (1 << (31 - n)))));
#define COUNT_1_INT(x) popcnt(x)
#define COUNT_0_INT(x) (32 - COUNT_1_INT(x))

#define MATRIX_OFF(x, y) (x * MATRIX_WIDTH + y)

#define OVERBOUND(x) (!!(x >> 3))
#define OVERFLOW(x, y) (OVERBOUND(x) | OVERBOUND(y))
//#define OFF_OVERBOUND(o) (OVERBOUND(o % 8) | OVERBOUND(o / 8))
//#define OFFVEC_OVERFLOW(o, d) (OFF_OVERBOUND(o) | OVERBOUND(o / 8 + DX(d)) | OVERBOUND(o % 8 + DY(d)))
//#define VECTOR_OVERFLOW(x, y, dx, dy) (OVERFLOW(x, y) | OVERBOUND(x + dx) | OVERBOUND(y + dy))

//p: panel, start of panel 0/1
#define GET_BIT_ONPANEL_BYOFF(b, p, o) GET_BIT_AT(b[2 * p +o / 32], o % 32)
#define GET_BIT_ONPANEL(b, p, x, y) GET_BIT_ONPANEL_BYOFF(b, p, MATRIX_OFF(x, y))
//#define GET_FACTION_BYOFF(b, o) (GET_BIT_ONPANEL_BYOFF(b, 0, o) + (GET_BIT_ONPANEL_BYOFF(b, 1, o) << 1))
#define GET_FACTION_AT(b, x, y) (GET_BIT_ONPANEL(b, 0, x, y) + (GET_BIT_ONPANEL(b, 1, x, y) << 1))
#define SET_BIT_ONPANEL_BYOFF(b, p, o, v) ;SET_BIT_AT(b[2 * p + o / 32], v, (o % 32));
#define SET_BIT_ONPANEL(b, p, x, y, v) SET_BIT_ONPANEL_BYOFF(b, p, MATRIX_OFF(x, y), v)
//#define SET_FACTION_BYOFF(b, o, f) SET_BIT_ONPANEL_BYOFF(b, 0, o, FV(f * 2));\
//                                   SET_BIT_ONPANEL_BYOFF(b, 1, o, FV(f * 2 + 1));
#define SET_FACTION_AT(b, x, y, f) SET_BIT_ONPANEL(b, 0, x, y, FV(f * 2));\
                                   SET_BIT_ONPANEL(b, 1, x, y, FV(f * 2 + 1));

//p: panel, 1 or 0
//#define PANEL_CNT(b, p) (COUNT_1_INT(b[2 * p]) + COUNT_1_INT(b[2 * p + 1]))
#define FACTION_CNT(b, f) (COUNT_1_INT(b[FOFF(f * 2)]) + COUNT_1_INT(b[FOFF(f * 2 + 1)]))
//#define BOARD_AND_CNT(b) (COUNT_1_INT(b[0] & b[2]) + COUNT_1_INT(b[1] & b[3]))
//#define BOARD_OR_CNT(b) (COUNT_1_INT(b[0] | b[2]) + COUNT_1_INT(b[1] | b[3]))

#define EMPTYSLOTS_ONBOARD(b) (COUNT_0_INT(b[0] | b[2]) + COUNT_0_INT(b[1] | b[3]))

//#define GET_ADJACENT_BIT_ONPANEL(b, x, y, d)

#define THREAD_NUM 1

int nextInt(int cap) {
    return cap == 1 ? 1 : 1 + rand() % (cap - 1);
}

std::ofstream fout;
int fileIndex = 0;
int recordCount = 0;
std::string filename_prefix = "/Users/sagnitude/log/record_";

pthread_mutex_t sum_mutex;

const int directionX[9] = {-1, 1,  0, 0, -1, -1,  1, 1, 0};

const int directionY[9] = { 0, 0, -1, 1, -1,  1, -1, 1, 0};

const int factionOffset[8] = {0, 0, 0, 2, 1, 3, 0, 0};

const int factionValues[8] = {0, 0, 1, 0, 0, 1, 1, 1};

const Faction factions[2] = {F1, F2};

const int fv[2] = {1, 2};

/**
 * Data Handler
 */
void applyPiecePlacementAtDirection(Board board, int x, int y, Faction f, Direction d, int count) {
    while (count != 0) {
        x += DX(d);
        y += DY(d);
        SET_FACTION_AT(board, x, y, f);
        count--;
    };
}

int getNext(Board board, int x, int y, Direction d) {
    if (OVERFLOW(x, y)) {
        return 0;
    }

    x += DX(d);
    y += DY(d);

    if (OVERFLOW(x, y)) {
        return 0;
    }

    return GET_FACTION_AT(board, x, y);
}

int countDirectionalFlippable(Board board, int x, int y, Faction f, Direction d) {
    int count = 0;
    int state = 0;
    while (1) {
        state = getNext(board, x, y, d);

        if (state == 0) {
            return 0;
        } else if (state == f) {
            return count;
        }

        x += DX(d);
        y += DY(d);

        count++;
    };
}

int countFlippable(Board board, int x, int y, Faction f) {
    int count = countDirectionalFlippable(board, x, y, f, TOP);
    count += countDirectionalFlippable(board, x, y, f, DOWN);
    count += countDirectionalFlippable(board, x, y, f, LEFT);
    count += countDirectionalFlippable(board, x, y, f, RIGHT);
    count += countDirectionalFlippable(board, x, y, f, TOPLEFT);
    count += countDirectionalFlippable(board, x, y, f, TOPRIGHT);
    count += countDirectionalFlippable(board, x, y, f, DOWNLEFT);
    count += countDirectionalFlippable(board, x, y, f, DOWNRIGHT);
    return count;
}

int countAndListFlippable(Board board, int x, int y, Faction f, int* flippable) {
    int count = flippable[0] = countDirectionalFlippable(board, x, y, f, TOP);
    count += flippable[1] = countDirectionalFlippable(board, x, y, f, DOWN);
    count += flippable[2] = countDirectionalFlippable(board, x, y, f, LEFT);
    count += flippable[3] = countDirectionalFlippable(board, x, y, f, RIGHT);
    count += flippable[4] = countDirectionalFlippable(board, x, y, f, TOPLEFT);
    count += flippable[5] = countDirectionalFlippable(board, x, y, f, TOPRIGHT);
    count += flippable[6] = countDirectionalFlippable(board, x, y, f, DOWNLEFT);
    count += flippable[7] = countDirectionalFlippable(board, x, y, f, DOWNRIGHT);
    return count;
}

int listAllCapablePoints(Board board, Faction f, int* capable) {
    int f1 = board[0] | board[2];
    int f2 = board[1] | board[3];
    int count = COUNT_0_INT(f1) + COUNT_0_INT(f2);
    if (count == 0) {
        return 0;
    }
    int c = 0;
    capable[c++] = 0;
    for (int i = 0; i < 32; i++) {
        if (GET_BIT_AT(f1, i) == 0 && countFlippable(board, i / MATRIX_WIDTH, i % MATRIX_WIDTH, f) != 0) {
            capable[c++] = i;
        }
        if (GET_BIT_AT(f2, i) == 0 && countFlippable(board, 4 + i / MATRIX_WIDTH, i % MATRIX_WIDTH, f) != 0) {
            capable[c++] = 32 + i;
        }
    }
    return c;
}

int getNthCapable(Game game, Faction f, int n) {
    for (int i = 0; i < MATRIX_WIDTH; i++) {
        for (int j = 0; j < MATRIX_WIDTH; j++) {
            int onf = GET_FACTION_AT(game.capable, i, j);
            if (onf == 3 || onf == f) {
                n--;
            }
            if (n == 0) {
                return MATRIX_OFF(i, j);
            }
        }
    }
    return -1;
}

int countAllCapablePoints(Board board, Faction f, int *cache) {
    int count = listAllCapablePoints(board, f, cache);
    return count;
}

/**
 * Player
 */
int randomSelectBot(Board board, Faction f, int* cache) {
    int count = listAllCapablePoints(board, f, cache);
    if (count == 1) {
        return ~0;
    }
    int c = cache[nextInt(count)];
    LOG(
            std::cout << "Capable points: " << (count) << " ";
    )
    return c;
}

int randomGameSelectorBot(Game game, Faction f, int* cache) {
    int c = f == F1 ? game.c1 : game.c2;
    int cap = getNthCapable(game, f, nextInt(c));
    LOG(
//            std::cout << "Capable points: " << (c) << " ";
    )
    return cap;
}

/**
 * Debug
 */
void printBoard(Board board) {
    std::cout << std::endl;
    for (int i = 0; i < MATRIX_WIDTH; i++) {
        for (int j = 0; j < MATRIX_WIDTH; j++) {
            int f = GET_FACTION_AT(board, i, j);
            if (f == 0) {
                std::cout << "- ";
            } else {
                std::cout << f << " ";
            }
        }
        std::cout << std::endl;
    }
}

int endGameCheck(Game game) {
    //board full?
    int empty = EMPTYSLOTS_ONBOARD(game.board);
    if (empty == 0) {
        LOG(
                std::cout << " board full, end game " << std::endl;
        )
        return 1;
    }
    //any faction empty?
    int f1 = game.f1;
    int f2 = game.f2;
    if (f1 == 0 || f2 == 0) {
        LOG(
                std::cout << " someone failed, end game " << std::endl;
        )
        return 1;
    }
    //no one has a capable point?
    int c1 = countAllCapablePoints(game.board, factions[0], game.capableCache);
    int c2 = countAllCapablePoints(game.board, factions[1], game.capableCache);
    LOG(
            std::cout << " checking... real: (" << c1 << ", " << c2 << ") estimated: (" << game.c1 << ", " << game.c2 << ")" << std::endl;
    )
    if (c1 == 0 && c2 == 0) {
        LOG(
                std::cout << " no one has a proper point, end game " << std::endl;
        )
        return 1;
    }
    if (game.steps > 70) {
        LOG(
                std::cout << " game over 128 steps, game failed " << std::endl;
        )
        return 1;
    }
    LOG(
            std::cout << " game goes on " << std::endl;
    )
    return 0;
}

void endGame(int* records, int count, Faction f) {
    LOG(
            std::cout << "Game end: with " << count << " steps" << std::endl;
    )

    for (int i = 0; i < count; i++) {
        int b = records[i * 9 + 8] >> 24;
        records[i * 9 + 8] |= (b == f);
    }

    REC(
            recordCount++;
            if (recordCount == 100000) {
                recordCount = 0;
                fileIndex++;
                fout.close();
                std::stringstream ss;
                ss << filename_prefix << fileIndex << ".log";
                fout.open(ss.str());
            }
            for (int i = 0; i < count; i++) {
                fout << records[i * 9 + 0] << "," << records[i * 9 + 1] << "," << records[i * 9 + 2] << "," << records[i * 9 + 3] << "," << records[i * 9 + 4];
                fout << "," << records[i * 9 + 5] << "," << records[i * 9 + 6] << "," << records[i * 9 + 7] << "," << records[i * 9 + 8] << std::endl;
            }
    )
}

int startGame(Game game, int (*p1)(Board, Faction, int*), int(*p2)(Board, Faction, int*)) {
    game.board[0] = 24;
    game.board[1] = 0;
    game.board[2] = 0;
    game.board[3] = 402653184;
    game.board[4] = 0;
    game.players[0] = p1;
    game.players[1] = p2;

    int x = 0;
    int y = 0;
    Faction currentFaction;
    int *flippables = (int *) malloc (8 * sizeof(int));
    int flippableCount = 0;

    /*
     * Calculate Inner Capable Matrix
     */
    while(true) {
        game.step++;
        game.step %= 2;

        //calculate capable points
        for (int i = 0; i < MATRIX_WIDTH; i++) {
            for (int j = 0; j < MATRIX_WIDTH; j++) {
                int bit = GET_FACTION_AT(game.board, i, j);
                int f = None;
                if (bit == 0) {
                    //inspect
                    int c1 = countFlippable(game.board, i, j, F1);
                    int c2 = countFlippable(game.board, i, j, F2);
                    f = c1 + c2 * 2;
                    if (c1 != 0) {
                        game.c1 ++;
                    }
                    if (c2 != 0) {
                        game.c2 ++;
                    }
                }
                SET_FACTION_AT(game.capable, i, j, (Faction) f);
            }
        }

        game.f1 = FACTION_CNT(game.board, F1);
        game.f2 = FACTION_CNT(game.board, F2);

//        game.action = randomGameSelectorBot(game, factions[game.step]);
        game.action = game.players[game.step](game.board, factions[game.step], game.capableCache);

        x = game.action / MATRIX_WIDTH;
        y = game.action % MATRIX_WIDTH;
        currentFaction = factions[game.step];

        if (game.action == 0) {
            std::cout << std::endl;
        }

        LOG(
                std::cout << "Step " << game.steps << ", player " << game.step << " at (" << game.action / MATRIX_WIDTH << ", " << game.action % MATRIX_WIDTH << ") with offset " << game.action << std::endl;
        )

        /*
         * Generate record
         */
        //snapshot board and selection
        game.records[game.steps * 9 + 0] = game.board[0];
        game.records[game.steps * 9 + 1] = game.board[1];
        game.records[game.steps * 9 + 2] = game.board[2];
        game.records[game.steps * 9 + 3] = game.board[3];
        game.records[game.steps * 9 + 4] = game.capable[0];
        game.records[game.steps * 9 + 5] = game.capable[1];
        game.records[game.steps * 9 + 6] = game.capable[2];
        game.records[game.steps * 9 + 7] = game.capable[3];
        int r = (fv[game.step] & 0xFF) << 24;
        r |= (game.action / MATRIX_WIDTH & 0xFF) << 16;
        r |= (game.action % MATRIX_WIDTH & 0xFF) << 8;
        game.records[game.steps * 9 + 8] = r;

        LOG(
                std::cout << "Capable matrix: " << std::endl;
                printBoard(game.capable);
        )

        /*
         * Put a piece at (x, y)
         */
        flippableCount = countAndListFlippable(game.board, x, y, currentFaction, flippables);
        if (flippableCount == 0) {
            continue;
        }
        int faction = GET_FACTION_AT(game.board, x, y);
        if (faction != 0) {
            continue;
        }
        SET_FACTION_AT(game.board, x, y, currentFaction);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, TOP,       flippables[0]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, DOWN,      flippables[1]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, LEFT,      flippables[2]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, RIGHT,     flippables[3]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, TOPLEFT,   flippables[4]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, TOPRIGHT,  flippables[5]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, DOWNLEFT,  flippables[6]);
        applyPiecePlacementAtDirection(game.board, x, y, currentFaction, DOWNRIGHT, flippables[7]);

        LOG(
                std::cout << std::endl;
                std::cout << "Game Matrix: " << std::endl;
                printBoard(game.board);
        )

        game.steps++;

        /*
         * End game check
         */
        if (endGameCheck(game) == 1) {
            game.f1 = FACTION_CNT(game.board, F1);
            game.f2 = FACTION_CNT(game.board, F2);
            endGame(game.records, game.steps, game.f1 > game.f2 ? F1 : game.f1 == game.f2 ? None : F2);
            LOGR(
                    std::cout << std::endl;
                    std::cout << "Game End: " << std::endl;
                    printBoard(game.board);
            )
            break;
        }
    };
    free(flippables);
    return 1;
}

void *singleGame(void *args) {
    for (int i = 0; i < 1; i++) {
        LOG(
                std::cout << " running " << i << std::endl;
        )
        Game* g = (Game *) malloc (sizeof (struct Game));
        *g = Game();
        startGame(*g, randomSelectBot, randomSelectBot);
        free(g);
    }
    return nullptr;
}


/**
* CUDA Body
*/

static void CheckCudaErrorAux(const char *, unsigned, const char *, cudaError_t);
#define CUDA_CHECK_RETURN(value) CheckCudaErrorAux(__FILE__,__LINE__, #value, value)

/**
* CUDA kernel that computes reciprocal values for a given vector
*/
__global__ void reciprocalKernel(float *data, unsigned vectorSize) {
	unsigned idx = blockIdx.x*blockDim.x + threadIdx.x;
	if (idx < vectorSize)
		data[idx] = 1.0 / data[idx];
}

/**
* Host function that copies the data and launches the work on GPU
*/
float *gpuReciprocal(float *data, unsigned size)
{
	float *rc = new float[size];
	float *gpuData;

	CUDA_CHECK_RETURN(cudaMalloc((void **)&gpuData, sizeof(float)*size));
	CUDA_CHECK_RETURN(cudaMemcpy(gpuData, data, sizeof(float)*size, cudaMemcpyHostToDevice));

	static const int BLOCK_SIZE = 256;
	const int blockCount = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;
	reciprocalKernel << <blockCount, BLOCK_SIZE >> > (gpuData, size);

	CUDA_CHECK_RETURN(cudaMemcpy(rc, gpuData, sizeof(float)*size, cudaMemcpyDeviceToHost));
	CUDA_CHECK_RETURN(cudaFree(gpuData));
	return rc;
}

int *gpuCopyInts(int* data, unsigned size) {
	int *rc = new int[size];
	int *gpuData;

	CUDA_CHECK_RETURN(cudaMalloc((void **)&gpuData, sizeof(int)*size));
	CUDA_CHECK_RETURN(cudaMemcpy(gpuData, data, sizeof(int)*size, cudaMemcpyHostToDevice));

	static const int BLOCK_SIZE = 256;
	const int blockCount = (size + BLOCK_SIZE - 1) / BLOCK_SIZE;
	reciprocalKernel << <blockCount, BLOCK_SIZE >> > (gpuData, size);

	CUDA_CHECK_RETURN(cudaMemcpy(rc, gpuData, sizeof(int)*size, cudaMemcpyDeviceToHost));
	CUDA_CHECK_RETURN(cudaFree(gpuData));
	return rc;
}

float *cpuReciprocal(float *data, unsigned size)
{
	float *rc = new float[size];
	for (unsigned cnt = 0; cnt < size; ++cnt) rc[cnt] = 1.0 / data[cnt];
	return rc;
}


void initialize(float *data, unsigned size)
{
	for (unsigned i = 0; i < size; ++i)
		data[i] = .5*(i + 1);
}

int main(void)
{

	srand((unsigned int) time(NULL));

	std::stringstream ss;
	ss << filename_prefix << fileIndex << ".log";
	fout.open(ss.str());

	pthread_attr_t attr;
	pthread_attr_init( &attr );
	pthread_attr_setdetachstate( &attr, PTHREAD_CREATE_JOINABLE );
	pthread_mutex_init( &sum_mutex, NULL );

	pthread_t tids[THREAD_NUM];
	for( int i = 0; i < THREAD_NUM; ++i ) {
		int ret = pthread_create( &tids[i], NULL, singleGame, NULL );
		if( ret != 0 ) {
			LOG(
					std::cout << "pthread_create error:error_code=" << ret << std::endl;
			)
		}
	}
	pthread_attr_destroy( &attr );

	void *status;
	for( int i = 0; i < THREAD_NUM; ++i ) {
		int ret = pthread_join( tids[i], &status );
		if( ret != 0 ) {
			LOG(
					std::cout << "pthread_join error:error_code=" << ret << std::endl;
			)
		}
	}

	pthread_mutex_destroy( &sum_mutex );

	pthread_exit( NULL );

	static const int WORK_SIZE = 65530;
	float *data = new float[WORK_SIZE];

	initialize(data, WORK_SIZE);

	float *recCpu = cpuReciprocal(data, WORK_SIZE);
	float *recGpu = gpuReciprocal(data, WORK_SIZE);
	float cpuSum = std::accumulate(recCpu, recCpu + WORK_SIZE, 0.0);
	float gpuSum = std::accumulate(recGpu, recGpu + WORK_SIZE, 0.0);

	/* Verify the results */
	std::cout << "gpuSum = " << gpuSum << " cpuSum = " << cpuSum << std::endl;

	/* Free memory */
	delete[] data;
	delete[] recCpu;
	delete[] recGpu;

	return 0;
}

/**
* Check the return value of the CUDA runtime API call and exit
* the application if the call has failed.
*/
static void CheckCudaErrorAux(const char *file, unsigned line, const char *statement, cudaError_t err)
{
	if (err == cudaSuccess)
		return;
	std::cerr << statement << " returned " << cudaGetErrorString(err) << "(" << err << ") at " << file << ":" << line << std::endl;
	exit(1);
}
